#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable

void AtomicAdd(volatile __global float *source, const float operand) {
	//union() function can be used to access bit values of a variable
    union { 
        unsigned int intVal;
        float floatVal;
    } newVal;
    union {
        unsigned int intVal;
        float floatVal;
    } prevVal;
    do {
        prevVal.floatVal = *source;
        newVal.floatVal = prevVal.floatVal + operand;
    } while (atomic_cmpxchg((volatile __global unsigned int *)source, prevVal.intVal, newVal.intVal) != prevVal.intVal);
}

__kernel
void coo_opencl(__global float *values, __global float *v, __global int *rowIdx, __global int *colIdx, const int nonZeros, __global float *answer) {
  const int __element = get_global_id(0);

  if(__element >= nonZeros)
	return;

	int i = __element;
	int _row = rowIdx[i];
	int _col = colIdx[i];
	float _value = values[i];
	AtomicAdd(&answer[_row], _value * v[_col]);
}

